#include<iostream>
#include<cuda_runtime.h>
#include<vector>
#include "../MyTimer.h"
#include<string>
#include<math.h>
#include<sstream>


void print_result(
    std::vector<std::string> &names, 
    std::vector<float>&times,
    std::vector<bool>& answer_right
)
{
    for(int i = 0 ; i < names.size(); i++)
    {
        std::cout<<names[i]<<" time:" << times[i] <<" result:"<<answer_right[i] <<std::endl;
    }
}
bool is_same(float*A, float*B, int n)
{
    for(int i = 0; i < n; i++)
        if(A[i] != B[i]) return 0;
    return 1;
}
void init(float* f, int n)
{
    for(int i = 0; i < n; i++)
        f[i] = rand()%100 + 1;
}
void common(float* A, float* B, float* C, int n)
{
    for(int i = 0; i < n; i++)
    {
        C[i] = A[i] + B[i];
    }
}
void common_run(
    float* A, float* B, float* answer, int n, 
    std::vector<std::string> &names, std::vector<float>&times, std::vector<bool>& answer_right
)
{
    GPUTimer t;
    t.start();
    common(A, B, answer, n);

    times.push_back(t.stop());
    names.push_back("无GPU");
    answer_right.push_back(1);
}


__global__ void GPU_common(float* A, float* B, float* C, int n, int repeat)
{
    int index = blockDim.x * blockIdx.x + threadIdx.x;
    //printf("%f, %f!\n");
    if(index < n)
        for(int i = 0; i < repeat; i++)
            C[index] = A[index] + B[index];
}
void GPU_run(
    float* A, float* B, float* answer, int n, 
    std::vector<std::string> &names, std::vector<float>&times, std::vector<bool>& answer_right,
    int type, int repeat, int nStreams
)
{

    float* c_A, *c_B, *c_C, *h_C;
    cudaHostAlloc((void**)&h_C, sizeof(float) * n, cudaHostAllocMapped);
    int nBs = sizeof(float) * n;
    cudaMalloc((void**)&c_A, nBs);
    cudaMalloc((void**)&c_B, nBs);
    cudaMalloc((void**)&c_C, nBs);


    int block_size = 1<<10;
    int grid_size = (n + block_size - 1) / block_size;
    GPUTimer t;
    t.start();

    if(type == 1)
    {
        cudaMemcpy(c_A, A, nBs, cudaMemcpyHostToDevice);
        cudaMemcpy(c_B, B, nBs, cudaMemcpyHostToDevice);
        GPU_common<<<grid_size, block_size>>>(c_A, c_B, c_C, n, repeat);
        cudaMemcpy(h_C, c_C, nBs, cudaMemcpyDeviceToHost);
        times.push_back(t.stop());
        names.push_back("有GPU无流");
    }
    else if(type == 2)
    {
        int streamSize = n / nStreams;
        int snBs = streamSize * sizeof(float);
        cudaStream_t streams[nStreams];
        for (int i = 0; i < nStreams; i++) {
            cudaStreamCreate(&streams[i]);
        }
        for(int i = 0; i < nStreams; i++)
        {
            int offset = i * streamSize;
            cudaMemcpyAsync(c_A + offset, A + offset, snBs, cudaMemcpyHostToDevice, streams[i]);
            cudaMemcpyAsync(c_B + offset, B + offset, snBs, cudaMemcpyHostToDevice, streams[i]);
            GPU_common<<<grid_size, block_size, 0, streams[i]>>>(c_A + offset, c_B + offset, c_C + offset, streamSize, repeat);
            cudaMemcpyAsync(h_C + offset, c_C + offset, snBs, cudaMemcpyDeviceToHost, streams[i]);//这个本质上只是一个提交。它不会卡住啥的，所以也能放在这里
        }
        for (int i = 0; i < nStreams; i++) {
            cudaStreamSynchronize(streams[i]);
        }

        times.push_back(t.stop());
                // 清理资源
        for (int i = 0; i < nStreams; i++) {
            cudaStreamDestroy(streams[i]);
        }

        names.push_back("有GPU有流");
    }


    answer_right.push_back(is_same(h_C, answer, n));
    


}


int main(int argc, char* argv[])
{
    int n = 1 << 20;
    int nStreams = 4;
    std::vector<std::string> names;
    std::vector<float> times;
    std::vector<bool> answer_right;

    float* A, *B, *answer;
    cudaHostAlloc((void**)&A, sizeof(float) * n, cudaHostAllocMapped);//cudaMemcpyAsync要求使用这个内存锁。
    cudaHostAlloc((void**)&B, sizeof(float) * n, cudaHostAllocMapped);
    cudaHostAlloc((void**)&answer, sizeof(float) * n, cudaHostAllocMapped);
    init(A, n);
    init(B, n);

    common_run(A, B, answer, n, names, times, answer_right);
    GPU_run(A, B, answer, n, names, times, answer_right, 1, 1, nStreams);
    GPU_run(A, B, answer, n, names, times, answer_right, 2, 1, nStreams);

    print_result(names, times, answer_right);
}





























